package ca.krakenpower.kvoxels;

/*
 * Based heavily on the work of Cyril Crassin.
 * http://maverick.inria.fr/Membres/Cyril.Crassin/thesis/CCrassinThesis_EN_Web.pdf
 * 
 * NOTES:
 * -sparse-twice voxel octree for empty/occluded space skipping
 * -brick maps
 * -pre-filter 3D mip-maps by trilinear interpolation of sub-nodes
 * -approximate cone tracing
 * -emission, absorption, scattering
 * -volume rendering integral (discretized) from s = s0 to s = D:
 * 
 * 		I(D) = SUM[i = 0, n: C_i * PRODUCT[j = i + 1, n: T_j]], C_0 = I(s_0)
 * 
 * 	where C represents the color contribution and T the transparency contribution.
 * 	Front-to-back evaluation can be represented recursively as:
 * 		
 * 		I(s_i) = I(s_(i + 1)) + T(s_(i + 1)) * C_i
 * 		T(s_i) = T(s_(i + 1)) * T_i
 * 	with
 * 		I(s_n) = C_n
 * 		T(s_n) = T_n
 * 	which leads to the following compositing scheme:
 * 		C_dst <- C_dst + (1 - a_dst) * C_src
 * 		a_dst <- a_dst + (1 - a_dst) * a_src
 * 
 * -geometry pre-filtering:
 * 		-compute average in-scattered energy (color?) and average transparency for a volume
 * 
 */

import static ca.krakenpower.kvoxels.Logger.log;

import java.awt.Font;
import java.io.BufferedReader;
import java.io.File;
import java.io.FileReader;
import java.io.IOException;
import java.net.URISyntaxException;
import java.nio.ByteBuffer;
import java.nio.FloatBuffer;
import java.nio.IntBuffer;
import java.util.List;
import java.util.concurrent.ThreadPoolExecutor;

import org.lwjgl.BufferUtils;
import org.lwjgl.LWJGLException;
import org.lwjgl.PointerBuffer;
import org.lwjgl.input.Keyboard;
import org.lwjgl.opencl.CL;
import org.lwjgl.opencl.CL10;
import org.lwjgl.opencl.CL10GL;
import org.lwjgl.opencl.CLCommandQueue;
import org.lwjgl.opencl.CLContext;
import org.lwjgl.opencl.CLDevice;
import org.lwjgl.opencl.CLEvent;
import org.lwjgl.opencl.CLKernel;
import org.lwjgl.opencl.CLMem;
import org.lwjgl.opencl.CLPlatform;
import org.lwjgl.opencl.CLProgram;
import org.lwjgl.opencl.OpenCLException;
import org.lwjgl.opencl.Util;
import org.lwjgl.opengl.Display;
import org.lwjgl.opengl.DisplayMode;
import org.lwjgl.opengl.GL11;
import org.lwjgl.opengl.GL12;
import org.lwjgl.opengl.GL15;
import org.lwjgl.opengl.GL21;
import org.lwjgl.opengl.GL30;
import org.newdawn.slick.Color;
import org.newdawn.slick.TrueTypeFont;

public class KVoxels {
	
	public static final String VERSION = "a012";
	
	private KVoxelOctree kvsd;
	
	private CLPlatform clPlatform;
	private List<CLDevice> clDevices;
	private CLContext clContext;
	private CLCommandQueue clCommandQueue;
	private CLMem clglRenderMem, nodeMem, dataMem, brickMem, vectorMem, renderParamsMem;
	private CLMem debugMem;
	private CLProgram rayProgram;
	private CLKernel rayKernel;
	
	private int globalWorkDimensions;
	private PointerBuffer globalWorkSize;
	private PointerBuffer localWorkSize;
	
	private ByteBuffer imageBuf;
	private FloatBuffer vectorParamBuf;
	private PointerBuffer imgOrigin, imgRegion, rp, sp;
	
	// PROFILING INFO
	private long kernelStart, kernelEnd;
	private long kernelWorkgroupSize;
	private PointerBuffer writeEvent, unmapEvent, kernelEvent;
	private CLEvent clevent;
	
	private Camera camera;
	
	private TrueTypeFont font;

	private int imageBufID;
	
	public void init() {
		
		Logger.initLogger(Logger.Type.STDOUT);
		
		try {
			
			initDisplay();
			
			initGL();
			
			initCL();
			
		} catch (LWJGLException e) {
			log("exception: " + e.getMessage());
			e.printStackTrace();
			closeOCL();
		} catch (OpenCLException e) {
			if (rayProgram != null) {
				log("CL_PROGRAM_BUILD_STATUS: " + rayProgram.getBuildInfoString(clDevices.get(0), CL10.CL_PROGRAM_BUILD_STATUS));
				log("CL_PROGRAM_BUILD_LOG: " + rayProgram.getBuildInfoString(clDevices.get(0), CL10.CL_PROGRAM_BUILD_LOG));
			}
			log("exception: " + e.getMessage());
			e.printStackTrace();
			closeOCL();
		}
		
		
	}
	
	private void initDisplay() throws LWJGLException {
		
		log("creating display window of size " + camera.getImageWidth() + "x" + camera.getImageHeight());
		
		Display.setDisplayMode(new DisplayMode(camera.getImageWidth(), camera.getImageHeight()));
		Display.setTitle("kVoxels " + VERSION);
		Display.create();
		
	}
	
	private void initGL() {
		
		log("initializing OpenGL...");
				
		GL11.glEnable(GL11.GL_TEXTURE_2D);
		GL11.glShadeModel(GL11.GL_SMOOTH);
		GL11.glDisable(GL11.GL_DEPTH_TEST);
		GL11.glDisable(GL11.GL_LIGHTING);
		
		GL11.glClearColor(0.0f, 0.0f, 0.0f, 0.0f);
		GL11.glClearDepth(1);
		 
		GL11.glBlendFunc(GL11.GL_SRC_ALPHA, GL11.GL_ONE_MINUS_SRC_ALPHA);
		 
		GL11.glViewport(0, 0, camera.getImageWidth(), camera.getImageHeight());
		GL11.glMatrixMode(GL11.GL_MODELVIEW);
		
		GL11.glMatrixMode(GL11.GL_PROJECTION);
		GL11.glLoadIdentity();
		GL11.glOrtho(0, camera.getImageWidth(), camera.getImageHeight(), 0, 1, -1);
		GL11.glMatrixMode(GL11.GL_MODELVIEW);
		
	}
	
	private void initCL() throws LWJGLException {
		
		writeEvent = BufferUtils.createPointerBuffer(1);
		unmapEvent = BufferUtils.createPointerBuffer(1);
		kernelEvent = BufferUtils.createPointerBuffer(1);
		
		log("initializing OpenCL...");
		IntBuffer errorBuff = BufferUtils.createIntBuffer(1);
		
		CL.create();
		
		clPlatform = CLPlatform.getPlatforms().get(0);
		log("using compute platform: " + clPlatform.getInfoString(CL10.CL_PLATFORM_NAME));
		
		clDevices = clPlatform.getDevices(CL10.CL_DEVICE_TYPE_GPU);
		
		clContext = CLContext.create(clPlatform, clDevices, null, Display.getDrawable(), errorBuff);
		
		Util.checkCLError(errorBuff.get(0));
		log("using compute device: " + clDevices.get(0).getInfoString(CL10.CL_DEVICE_NAME));
		log("    global mem size: " + clDevices.get(0).getInfoInt(CL10.CL_DEVICE_GLOBAL_MEM_SIZE));
		log("    global cache size: " + clDevices.get(0).getInfoInt(CL10.CL_DEVICE_GLOBAL_MEM_CACHE_SIZE));
		log("    local mem size: " + clDevices.get(0).getInfoInt(CL10.CL_DEVICE_LOCAL_MEM_SIZE));
		log("    max constant args: " + clDevices.get(0).getInfoInt(CL10.CL_DEVICE_MAX_CONSTANT_ARGS));
		log("    max workgroup size: " + clDevices.get(0).getInfoInt(CL10.CL_DEVICE_MAX_WORK_GROUP_SIZE));
		log("    max workitem dimensions: " + clDevices.get(0).getInfoInt(CL10.CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS));
		
		clCommandQueue = CL10.clCreateCommandQueue(clContext, clDevices.get(0), CL10.CL_QUEUE_PROFILING_ENABLE, errorBuff);
		Util.checkCLError(errorBuff.get(0));
		
		// load the kernel
		rayProgram = CL10.clCreateProgramWithSource(clContext, loadText("castRay2.cl"), errorBuff);
		Util.checkCLError(errorBuff.get(0));
		
		int error = CL10.clBuildProgram(rayProgram, clDevices.get(0), "-Werror -cl-mad-enable -cl-fast-relaxed-math", null);
		Util.checkCLError(error);
								
		rayKernel = CL10.clCreateKernel(rayProgram, "castRay", errorBuff);
		Util.checkCLError(errorBuff.get(0));
	
		imageBuf = BufferUtils.createByteBuffer(camera.getImageWidth() * camera.getImageHeight() * 4);
		imageBuf.clear();
		
		ByteBuffer format = BufferUtils.createByteBuffer(8);
		format.putInt(CL10.CL_RGBA);
		format.putInt(CL10.CL_FLOAT);
		format.rewind();
		
		imageBufID = GL15.glGenBuffers();
		GL15.glBindBuffer(GL21.GL_PIXEL_UNPACK_BUFFER, imageBufID);
		GL15.glBufferData(GL21.GL_PIXEL_UNPACK_BUFFER, imageBuf, GL15.GL_STREAM_DRAW);
		
		clglRenderMem = CL10GL.clCreateFromGLBuffer(clContext, CL10.CL_MEM_WRITE_ONLY, imageBufID, errorBuff);
		
		Util.checkCLError(errorBuff.get(0));
		GL15.glBindBuffer(GL21.GL_PIXEL_UNPACK_BUFFER, 0);
		
		imgOrigin = BufferUtils.createPointerBuffer(3);
		imgRegion = BufferUtils.createPointerBuffer(3);
		imgOrigin.put(0);
		imgOrigin.put(0);
		imgOrigin.put(0);
		imgRegion.put(camera.getImageWidth());
		imgRegion.put(camera.getImageHeight());
		imgRegion.put(1);
		imgOrigin.rewind();
		imgRegion.rewind();
		rp = BufferUtils.createPointerBuffer(1);
		rp.put(camera.getImageWidth() * 4 * 4);
		rp.rewind();
		sp = BufferUtils.createPointerBuffer(1);
		sp.put(0);
		sp.rewind();
		
		int[] nodes = kvsd.getNodePool();
		int[] data = kvsd.getDataPool();
		int[] bricks = kvsd.getBrickPool();
		
		log("node pool size: " + nodes.length);
		log("data pool size: " + data.length);
		log(bricks.length / kvsd.tBrick3Size + " bricks, " + bricks.length + " voxels");
		long memsize = (nodes.length * 4) + (data.length * 4) + (bricks.length * 4);
		log("actual data size: " + getMemSizeString(memsize));
		
		IntBuffer nodeBuf = BufferUtils.createIntBuffer(nodes.length);
		nodeBuf.put(nodes);
		nodeBuf.rewind();
		
		IntBuffer dataBuf = BufferUtils.createIntBuffer(data.length);
		dataBuf.put(data);
		dataBuf.rewind();
		
		nodeMem = CL10.clCreateBuffer(clContext, 
				CL10.CL_MEM_READ_ONLY | CL10.CL_MEM_COPY_HOST_PTR, 
				nodeBuf, 
				errorBuff);
		Util.checkCLError(errorBuff.get(0));
		dataMem = CL10.clCreateBuffer(clContext, 
				CL10.CL_MEM_READ_ONLY | CL10.CL_MEM_COPY_HOST_PTR, 
				dataBuf, 
				errorBuff);
		Util.checkCLError(errorBuff.get(0));
		
		int tsizeX, tsizeY, tsizeZ;
		tsizeX = tsizeY = tsizeZ = (int) (Math.cbrt(bricks.length / kvsd.tBrick3Size) + 1) * kvsd.totalBSize;
		int bsizeX, bsizeY, bsizeZ;
		bsizeX = bsizeY = bsizeZ = kvsd.totalBSize;
		
		IntBuffer brickBuf = BufferUtils.createIntBuffer(tsizeX * tsizeY * tsizeZ);
		for (int k = 0; k < tsizeZ; k++) {
			for (int j = 0; j < tsizeY; j++) {
				for (int i = 0; i < tsizeX; i++) {
					// convert the index into the texture (t) into an index into the brick pool (b)
					int bx, by, bz, bw, b, t;
					t = i + (j * tsizeX) + (k * tsizeY * tsizeX);
					
					bx = i % bsizeX;
					by = j % bsizeY;
					bz = k % bsizeZ;
					bw = (i / bsizeX) 
							+ ((j / bsizeY) * (tsizeX / bsizeX)) 
							+ ((k / bsizeZ) * (tsizeY / bsizeY) * (tsizeX / bsizeX));
					
					b = bx + (by * bsizeX) + (bz * bsizeY * bsizeX) + (bw * bsizeZ * bsizeY * bsizeX);
					
					if (b < bricks.length) {
						brickBuf.put(t, bricks[b]);
					} else {
						brickBuf.put(t, 0);
					}
				}
			}
		}
		brickBuf.rewind();
		
		int brickTextureID = GL11.glGenTextures();
		GL11.glBindTexture(GL12.GL_TEXTURE_3D, brickTextureID);
		
		GL12.glTexImage3D(GL12.GL_TEXTURE_3D,
				0,
				GL30.GL_RGBA8UI,
				tsizeX, tsizeY, tsizeZ, 
				0,
				GL30.GL_RGBA_INTEGER,
				GL11.GL_UNSIGNED_BYTE,
				brickBuf);
		
		brickMem = CL10GL.clCreateFromGLTexture3D(
				clContext, 
				CL10.CL_MEM_READ_ONLY,
				GL12.GL_TEXTURE_3D,
				0,
				brickTextureID,
				errorBuff);
		Util.checkCLError(errorBuff.get(0));
		
		vectorParamBuf = BufferUtils.createFloatBuffer(16);
		vectorParamBuf.rewind();
		vectorMem = CL10.clCreateBuffer(clContext,
				CL10.CL_MEM_READ_ONLY | CL10.CL_MEM_COPY_HOST_PTR,
				vectorParamBuf,
				errorBuff);
		
		FloatBuffer renderParamsBuf = BufferUtils.createFloatBuffer(8);
		renderParamsBuf.put((float) camera.fieldOfView());
		renderParamsBuf.put((float) kvsd.totalBSize);
		renderParamsBuf.put((float) kvsd.getSize());
		renderParamsBuf.put((float) camera.getNearPlaneSize()); // nearPlaneSize
		renderParamsBuf.put((float) camera.getNearPlaneDistance()); // nearPlaneDist
		renderParamsBuf.put((float) camera.getImageWidth());
		renderParamsBuf.put((float) camera.getImageHeight());
		renderParamsBuf.rewind();
		
		renderParamsMem = CL10.clCreateBuffer(clContext, 
				CL10.CL_MEM_READ_ONLY | CL10.CL_MEM_COPY_HOST_PTR, 
				renderParamsBuf, 
				errorBuff);
		Util.checkCLError(errorBuff.get(0));
		
		FloatBuffer debugBuf = BufferUtils.createFloatBuffer(16);
		debugMem = CL10.clCreateBuffer(clContext,
				CL10.CL_MEM_READ_WRITE,
				debugBuf,
				errorBuff);
		Util.checkCLError(errorBuff.get(0));
		
		globalWorkDimensions = 1;
		globalWorkSize = BufferUtils.createPointerBuffer(globalWorkDimensions);
		globalWorkSize.put(camera.getImageWidth() * camera.getImageHeight());
		globalWorkSize.rewind();
		localWorkSize = BufferUtils.createPointerBuffer(1);
		localWorkSize.put(128);
		localWorkSize.rewind();
		
		CL10GL.clEnqueueAcquireGLObjects(clCommandQueue, brickMem, null, null);
		Util.checkCLError(CL10.clFinish(clCommandQueue));
		
		rayKernel.setArg(0, nodeMem);
		rayKernel.setArg(1, dataMem);
		rayKernel.setArg(2, brickMem);
		rayKernel.setArg(3, vectorMem);
		rayKernel.setArg(4, renderParamsMem);
		rayKernel.setArg(5, clglRenderMem);
		rayKernel.setArg(6, debugMem);
		
		CL10GL.clEnqueueReleaseGLObjects(clCommandQueue, brickMem, null, null);
		Util.checkCLError(CL10.clFinish(clCommandQueue));
		
		Util.checkCLError(errorBuff.get(0));

	}
	
	public static String loadText(String name) {
		// http://www.lwjgl.org/wiki/index.php?title=OpenCL_in_LWJGL
		if(!name.endsWith(".cl")) {
			name += ".cl";
		}
		BufferedReader br = null;
		String resultString = null;
		try {
			// Get the file containing the OpenCL kernel source code
			File clSourceFile = new File(KVoxels.class.getClassLoader().getResource(name).toURI());
			// Create a buffered file reader to read the source file
			br = new BufferedReader(new FileReader(clSourceFile));
			// Read the file's source code line by line and store it in a string buffer
			String line = null;
			StringBuilder result = new StringBuilder();
			while((line = br.readLine()) != null) {
				result.append(line);
				result.append("\n");
			}
			// Convert the string builder into a string containing the source code to return
			resultString = result.toString();
		} catch(NullPointerException npe) {
			// If there is an error finding the file
			System.err.println("Error retrieving OpenCL source file: ");
			npe.printStackTrace();
		} catch(URISyntaxException urie) {
			// If there is an error converting the file name into a URI
			System.err.println("Error converting file name into URI: ");
			urie.printStackTrace();
		} catch(IOException ioe) {
			// If there is an IO error while reading the file
			System.err.println("Error reading OpenCL source file: ");
			ioe.printStackTrace();
		} finally {
			// Finally clean up any open resources
			try {
				br.close();
			} catch (IOException ex) {
				// If there is an error closing the file after we are done with it
				System.err.println("Error closing OpenCL source file");
				ex.printStackTrace();
			}
		}

		// Return the string read from the OpenCL kernel source code file
		return resultString;
	}
	
	private void closeOCL() {

		CL10.clReleaseMemObject(clglRenderMem);
		CL10.clReleaseMemObject(brickMem);
		CL10.clReleaseMemObject(dataMem);
		CL10.clReleaseMemObject(nodeMem);
		CL10.clReleaseMemObject(vectorMem);
		CL10.clReleaseMemObject(renderParamsMem);
		CL10.clReleaseMemObject(debugMem);
		
		CL10.clReleaseKernel(rayKernel);
		CL10.clReleaseProgram(rayProgram);
		CL10.clReleaseCommandQueue(clCommandQueue);
		CL10.clReleaseContext(clContext);
		
		CL.destroy();
		
	}
	
	private void renderCL() {

		vectorParamBuf.put(0, (float) camera.localForward().x);
		vectorParamBuf.put(1, (float) camera.localForward().y);
		vectorParamBuf.put(2, (float) camera.localForward().z);
		vectorParamBuf.put(3, (float) 0);

		vectorParamBuf.put(4, (float) camera.localLeft().x);
		vectorParamBuf.put(5, (float) camera.localLeft().y);
		vectorParamBuf.put(6, (float) camera.localLeft().z);
		vectorParamBuf.put(7, (float) 0);

		vectorParamBuf.put(8, (float) camera.localDown().x);
		vectorParamBuf.put(9, (float) camera.localDown().y);
		vectorParamBuf.put(10, (float) camera.localDown().z);
		vectorParamBuf.put(11, (float) 0);

		vectorParamBuf.put(12, (float) camera.getPosition().x);
		vectorParamBuf.put(13, (float) camera.getPosition().y);
		vectorParamBuf.put(14, (float) camera.getPosition().z);
		vectorParamBuf.put(15, (float) 0);

//		float[] debug = new float[16];
		
		try {
			
			Util.checkCLError(CL10.clEnqueueWriteBuffer(clCommandQueue, vectorMem, CL10.CL_TRUE, 0, vectorParamBuf, null, writeEvent));
			
			CL10GL.clEnqueueAcquireGLObjects(clCommandQueue, clglRenderMem, null, null);
			
			Util.checkCLError(CL10.clEnqueueNDRangeKernel(clCommandQueue, rayKernel, globalWorkDimensions, null, globalWorkSize, localWorkSize, writeEvent, kernelEvent));
			
			CL10GL.clEnqueueReleaseGLObjects(clCommandQueue, clglRenderMem, null, null);
						
			clevent = clCommandQueue.getCLEvent(kernelEvent.get(0));
			CL10.clWaitForEvents(clevent);
			kernelStart = clevent.getProfilingInfoLong(CL10.CL_PROFILING_COMMAND_QUEUED);
			kernelEnd = clevent.getProfilingInfoLong(CL10.CL_PROFILING_COMMAND_END);
			
//
//			FloatBuffer debugBuf = BufferUtils.createFloatBuffer(16);
//			CL10.clEnqueueReadBuffer(clCommandQueue, debugMem, CL10.CL_TRUE, 0, debugBuf, null, null);
//			CL10.clFinish(clCommandQueue);
//			debugBuf.rewind();
//			debugBuf.get(debug);
			
		} catch (Exception e) {
			log("render exception: " + e.getMessage());
			e.printStackTrace();
			closeOCL();
		}
		
//		GL15.glBindBuffer(GL21.GL_PIXEL_UNPACK_BUFFER, imageBufID);
//		
//		GL11.glDrawPixels(camera.getImageWidth(), 
//				camera.getImageHeight(), 
//				GL12.GL_BGRA,
//				GL12.GL_UNSIGNED_INT_8_8_8_8, 
//				0);
//		
//		GL15.glBindBuffer(GL21.GL_PIXEL_UNPACK_BUFFER, 0);
			
	}
	
	private void renderGL() {
		
		GL11.glClear(GL11.GL_COLOR_BUFFER_BIT);
		
		GL15.glBindBuffer(GL21.GL_PIXEL_UNPACK_BUFFER, imageBufID);
		
		GL11.glDrawPixels(camera.getImageWidth(), 
				camera.getImageHeight(), 
				GL12.GL_BGRA,
				GL12.GL_UNSIGNED_INT_8_8_8_8, 
				0);
		
		GL15.glBindBuffer(GL21.GL_PIXEL_UNPACK_BUFFER, 0);
		
	}

	public void run() {
		
		Font awtFont = new Font(Font.MONOSPACED, 0, 12);
		font = new TrueTypeFont(awtFont, false);
		
		double moveSpeed = 5;
		double rotationSpeed = Math.PI / 128;
		
		Display.setVSyncEnabled(false);
		int frameCount = 1;
		long t1, t2, deltaT;
		double fps = 0;
		
		CL10GL.clEnqueueAcquireGLObjects(clCommandQueue, brickMem, null, null);
		
		t1 = System.nanoTime();
		while (!Display.isCloseRequested()) {
			
			
			if (Keyboard.isKeyDown(Keyboard.KEY_W))
				camera.move(moveSpeed, camera.localForward());
			if (Keyboard.isKeyDown(Keyboard.KEY_S))
				camera.move(moveSpeed, camera.localBackward());
			
			if (Keyboard.isKeyDown(Keyboard.KEY_D))
				camera.move(moveSpeed, camera.localRight());
			if (Keyboard.isKeyDown(Keyboard.KEY_A))
				camera.move(moveSpeed, camera.localLeft());
			
			if (Keyboard.isKeyDown(Keyboard.KEY_R))
				camera.move(moveSpeed, camera.localUp());
			if (Keyboard.isKeyDown(Keyboard.KEY_F))
				camera.move(moveSpeed, camera.localDown());
			
			if (Keyboard.isKeyDown(Keyboard.KEY_T))
				camera.rotate(rotationSpeed, camera.localLeft());
			if (Keyboard.isKeyDown(Keyboard.KEY_G))
				camera.rotate(rotationSpeed, camera.localRight());
			
			if (Keyboard.isKeyDown(Keyboard.KEY_C))
				camera.rotate(rotationSpeed, camera.localUp());
			if (Keyboard.isKeyDown(Keyboard.KEY_Z))
				camera.rotate(rotationSpeed, camera.localDown());

			if (Keyboard.isKeyDown(Keyboard.KEY_E))
				camera.rotate(rotationSpeed, camera.localForward());
			if (Keyboard.isKeyDown(Keyboard.KEY_Q))
				camera.rotate(rotationSpeed, camera.localBackward());
				
			
			
			// render OpenGL here
			
			renderCL();
			renderGL();
			
			frameCount++;
			if (frameCount % 5 == 0) {
				t2 = System.nanoTime();
				deltaT = t2 - t1;
				fps = frameCount / (deltaT / 1000000000.0);
				t1 = t2;
				frameCount = 0;
			}
			long runTimeNS;
			double runTimeMS;
			runTimeNS = kernelEnd - kernelStart;
			runTimeMS = runTimeNS / 1000000.0;
			kernelWorkgroupSize = rayKernel.getWorkGroupInfoSize(clDevices.get(0), CL10.CL_KERNEL_WORK_GROUP_SIZE);
			
			GL11.glEnable(GL11.GL_BLEND);
			font.drawString(10, 10, fps + " fps", Color.white);
			font.drawString(10, 25, "camera pos: " + camera.getPosition());
			font.drawString(10, 40, "local forward: " + camera.localForward());
			font.drawString(10, 55, "vsd size: " + kvsd.getSize());
			font.drawString(10, 70, "kernel run time: " + runTimeMS + " ms (" + runTimeNS + " ns)");
			font.drawString(10, 85, "unaccounted frame time: " + ((1000.0 / fps) - runTimeMS) + " ms");
			GL11.glDisable(GL11.GL_BLEND);

			Display.update();
//			Display.sync(60);
			
		}
		

		CL10GL.clEnqueueReleaseGLObjects(clCommandQueue, brickMem, null, null);
		
		Display.destroy();
		closeOCL();
	}
	
	public static void main(String[] args) {
		
//		System.setProperty("org.lwjgl.util.Debug", "true");
		
		Logger.initLogger(Logger.Type.STDOUT);
		
		int ngSize = 4;
		log("generating perlin noise function...");
		log("noise scale: " + ngSize);
		log("noise seed: " + 0);
		NoiseGen ng = new NoiseGen(ngSize, 255);
		ng.generateNoiseFunction();
		
		KVoxels kv = new KVoxels();
		
		int kvsdSize = 256;
		int voxelSkip = 1;
		int vb = 64;
		
		log("creating KVoxelOctree of size " + kvsdSize);
		kv.kvsd = new KVoxelOctree((int) kvsdSize, 16);
		log("maximum subdivisions: " + kv.kvsd.getMaxSubdivisions());
		
		log("filling tree with voxel data...");
		
		long t1, t2, dt;
		t1 = System.nanoTime();
		for (int k = vb; k < kvsdSize - vb; k += voxelSkip) {
			if (k % 32 == 0) {
				log(((float) k / kvsdSize * 100) + "%... (" + (k * kvsdSize * kvsdSize / voxelSkip) + " insertions)");
			}
			for (int j = vb; j < kvsdSize - vb; j += voxelSkip) {
				for (int i = vb; i < kvsdSize - vb; i += voxelSkip) {
								
//					if (i <= vb || j <= vb || k <= vb
//							|| i > kvsdSize - vb || j > kvsdSize - vb || k > kvsdSize - vb) {
//						kv.kvsd.insert(0xFF00, i, j, k);
//					} else {
						double r = (double) (ngSize - 1) / (double) (kvsdSize);
						int a = (int) (ng.getValue(Vector3.get((i) * r, (j) * r, (k) * r)) * 8);
						a = Math.abs(a & 8);
//						a = (a >= 4) ? 0xFF : 0x00;
						kv.kvsd.insert(packRGBA(0, 0x1f, 0x3f, a), i, j, k);
//					}
						
				}
			}
		}
		t2 = System.nanoTime();
		dt = t2 - t1;
		log((kvsdSize * kvsdSize * kvsdSize / voxelSkip) + " insertions completed in " + (dt / 1000000000.0) + " s");
		log((dt / (kvsdSize * kvsdSize * kvsdSize / voxelSkip)) + " ns per insertion (avg)");
		
		log("node pool expansions performed: " + kv.kvsd.getNodeExpansions());
		log("brick pool expansions performed: " + kv.kvsd.getBrickExpansions());
		log("brick chunking operations performed: " + kv.kvsd.getNumBrickChunks());
		log("chunking nodes...");
		kv.kvsd.chunkSubnodes(0);
		log("node chunking operations performed: " + kv.kvsd.getNumNodeChunks());
		
		kv.kvsd.shrinkPools();
		
		log("final node count: " + kv.kvsd.getNumNodes());
		log("final brick count: " + kv.kvsd.getNumBricks());
		log("final voxel count: " + kv.kvsd.getNumVoxels());
		
		long memSize = (kv.kvsd.getNumNodes() * 8 + kv.kvsd.getNumVoxels() * 4);
		log("estimated final size (in bytes): " + getMemSizeString(memSize));
	
		kv.camera = new Camera(1024, Math.PI * 1 / 3, 1);
		kv.camera.setImageSize(1024, 768);
		kv.camera.setPosition(Vector3.get(kvsdSize / 2, kvsdSize / 2, kvsdSize * -2));
		
		kv.init();
		kv.run();
	
	}
	
	public static String getMemSizeString(long memSize) {
		String memSizeStr;
		if (memSize > 8192L) {
			if (memSize > 8192L * 1024L) {
				if (memSize > 8192L * 1024L * 1024L) {
					memSizeStr = Long.toString(memSize / (1024L * 1024L * 1024L)) + " GB";
				} else {
					memSizeStr = Long.toString(memSize / (1024L * 1024L)) + " MB";
				}
			} else {
				memSizeStr = Long.toString(memSize / (1024L)) + " KB";
			}
		} else {
			memSizeStr = Long.toString(memSize) + " B";
		}
		
		return memSizeStr;
	}
	
	public static int packRGBA(int r, int g, int b, int a) {
		int clr = a & 0xFF;
		clr |= (b & 0xFF) << 8;
		clr |= (g & 0xFF) << 16;
		clr |= (r & 0xFF) << 24;
		return clr;
	}

}
